#pragma once
#include <geometry\primitive.h>
#include "lbvh_structure.h"


namespace gc {
	__global__ void buildBV_SOA(int size, BV_SOA bvs, const SimpleSphere *prims);

	__global__ void buildBVs(int size, LBvhLeaves leaves, const SimpleSphere *prims);

	///< parallel reduce in O(logn). Stage 1
	__global__ void calcSceneBV_1(int size, LBvhLeaves leaves, BV* d_out);

	///< Stage 2: Keep calling kernel until size is reduced to 1.
	__global__ void calcSceneBV_2(int size, BV* d_in, BV* d_out);

	__global__ void calcMortonCodes_32b(int size, LBvhLeaves leaves, BV* scene);

	__global__ void reorderLeaves(int size, LBvhLeaves leaves, LBvhLeaves leaves_sorted, int *mapidx);

	__global__ void calcBuildMetrics(int interiorSize, LBvhLeaves leaves, uint *metrics);

	__global__ void buildInteriors(int leafSize, LBvhLeaves leaves, LBvhInteriors interiors, InteriorBuildAid aid);

	__global__ void calcLeafEscape(int size, LBvhLeaves leaves, LBvhInteriors interiors);

	__global__ void calcInteriorEscape(int interiorSize, LBvhLeaves leaves, LBvhInteriors interiors);

	///< histogram sort
	__global__ void calcInteriorNewIdx(int size, LBvhLeaves leaves, LBvhInteriors interiors, int *leftLeafCount, int *mapidx);

	__global__ void reorderInteriors(int interiorSize, LBvhLeaves leaves, LBvhInteriors interiors, LBvhInteriors interiors_sorted, int *mapidx);

	__global__ void resetBVs(int size, AABB_SOA bv_soa);

}